#define K_MAX_SHAPE_DIM 0
#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
class KernelLcm {
public:
    __aicore__ inline KernelLcm() {}
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR other, GM_ADDR out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        if (GetBlockIdx() < bigNum) {
            coreDataNum = bigLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + bigLength * GetBlockIdx(), bigLength);
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + bigLength * GetBlockIdx(), bigLength);
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + bigLength * GetBlockIdx(), bigLength);
        } else {
            coreDataNum = smallLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + bigLength * bigNum + smallLength * (GetBlockIdx() - bigNum), smallLength);
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + bigLength * bigNum + smallLength * (GetBlockIdx() - bigNum), smallLength);
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + bigLength * bigNum + smallLength * (GetBlockIdx() - bigNum), smallLength);
        }
        this->tileDataLenght = 5632;
        this->loopCount = coreDataNum/this->tileDataLenght;
        this->tailDataLenght = coreDataNum - this->tileDataLenght*this->loopCount;
        if(this->tailDataLenght == 0){this->tailDataLenght = this->tileDataLenght; this->loopCount--;}
        pipeIn->InitBuffer(inQueueInput, BUFFER_NUM, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(inQueueOther, BUFFER_NUM, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(QueueTmp1, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(QueueTmp2, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(QueueTmpMask, this->tileDataLenght/8 * sizeof(uint8_t));
    }
    __aicore__ inline void Process() {
        this->processDataLenght = this->tileDataLenght;
        for(int32_t i=0; i<this->loopCount; i++)
        {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
        this->processDataLenght = this->tailDataLenght;
        CopyIn(this->loopCount);
        Compute(this->loopCount);
        CopyOut(this->loopCount);
    }
private:
    __aicore__ inline void CopyIn(int32_t progress) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.AllocTensor<DTYPE_INPUT>();
        LocalTensor<DTYPE_INPUT> otherLocal = inQueueOther.AllocTensor<DTYPE_INPUT>();
        DataCopy(inputLocal, inputGm[progress * this->tileDataLenght], this->processDataLenght);
        DataCopy(otherLocal, otherGm[progress * this->tileDataLenght], this->processDataLenght);
        inQueueInput.EnQue(inputLocal);
        inQueueOther.EnQue(otherLocal);
    }
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.DeQue<DTYPE_INPUT>();
        LocalTensor<DTYPE_INPUT> otherLocal = inQueueOther.DeQue<DTYPE_INPUT>();
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.AllocTensor<DTYPE_INPUT>();
        if constexpr (std::is_same_v<DTYPE_INPUT, int32_t>) 
        {
            auto tmp1 = QueueTmp1.Get<float>();
            auto tmp2 = QueueTmp2.Get<float>();
            auto maskLocal = QueueTmpMask.Get<uint8_t>();
            Cast(inputLocal.ReinterpretCast<float>(), inputLocal, RoundMode::CAST_NONE, this->processDataLenght);
            Cast(otherLocal.ReinterpretCast<float>(), otherLocal, RoundMode::CAST_NONE, this->processDataLenght);
            Abs(inputLocal.ReinterpretCast<float>(), inputLocal.ReinterpretCast<float>(), this->processDataLenght);
            Abs(otherLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), this->processDataLenght);
            Mul(outLocal.ReinterpretCast<float>(), inputLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), this->processDataLenght);
            Duplicate(tmp2, (float)1, this->processDataLenght);
            for(int i = 0; i<10; i++)
            {
                Max(tmp1, inputLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), this->processDataLenght);
                Min(inputLocal.ReinterpretCast<float>(), inputLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), this->processDataLenght);
                Div(otherLocal.ReinterpretCast<float>(), tmp1, inputLocal.ReinterpretCast<float>(), this->processDataLenght);
                Cast(otherLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), RoundMode::CAST_FLOOR, this->processDataLenght);
                Mul(otherLocal.ReinterpretCast<float>(), otherLocal.ReinterpretCast<float>(), inputLocal.ReinterpretCast<float>(), this->processDataLenght);
                Sub(otherLocal.ReinterpretCast<float>(), tmp1, otherLocal.ReinterpretCast<float>(), this->processDataLenght);
                CompareScalar(maskLocal, otherLocal.ReinterpretCast<float>(), (float)0, CMPMODE::EQ, this->tileDataLenght);
                Select(tmp2, maskLocal, inputLocal.ReinterpretCast<float>(), tmp2, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataLenght);
            }
            Div(outLocal.ReinterpretCast<float>(), outLocal.ReinterpretCast<float>(), tmp2, this->processDataLenght);
            Cast(outLocal, outLocal.ReinterpretCast<float>(), RoundMode::CAST_RINT, this->processDataLenght);
        }
        else
        {
            for(int32_t i=0; i<this->processDataLenght; i++)
            {
                outLocal.SetValue(i, lcm(inputLocal.GetValue(i), otherLocal.GetValue(i)));
            }
        }
        inQueueInput.FreeTensor(inputLocal);
        inQueueOther.FreeTensor(otherLocal);
        outQueueOut.EnQue<DTYPE_INPUT>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.DeQue<DTYPE_INPUT>();
        DataCopy(outGm[progress * this->tileDataLenght], outLocal, this->processDataLenght);
        outQueueOut.FreeTensor(outLocal);
    }
    __aicore__ inline DTYPE_INPUT gcd(DTYPE_INPUT m, DTYPE_INPUT n)
    {
        while(n != 0)
        {
            DTYPE_INPUT temp = m % n;
            m = n;
            n = temp;
        }
        return (DTYPE_INPUT)(m);
    }
    __aicore__ inline DTYPE_INPUT lcm(DTYPE_INPUT m, DTYPE_INPUT n)
    {
        int64_t m1 = m;
        int64_t n1 = n;
        while(n1 != 0)
        {
            DTYPE_INPUT temp = m1 % n1;
            m1 = n1;
            n1 = temp;
        }
        int64_t mn = m / m1 * n;
        if(mn < 0) mn = -mn;
        return (DTYPE_INPUT)(mn);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInput, inQueueOther;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmpMask;
    GlobalTensor<DTYPE_INPUT> inputGm;
    GlobalTensor<DTYPE_INPUT> otherGm;
    GlobalTensor<DTYPE_INPUT> outGm;
    uint32_t loopCount;
    uint32_t tileDataLenght;
    uint32_t tailDataLenght;
    uint32_t processDataLenght;
};
class KernelLcmBC {
public:
    __aicore__ inline KernelLcmBC() {}
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR other, GM_ADDR out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, 
                                uint32_t repeat, uint32_t cycles, uint32_t interval, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        if (GetBlockIdx() < bigNum) {
            coreDataNum = bigLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + bigLength * cycles * interval * GetBlockIdx());
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + bigLength * interval * (GetBlockIdx()));
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + bigLength * cycles * interval * GetBlockIdx());
        } else {
            coreDataNum = smallLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + bigLength * cycles * interval  * bigNum + smallLength * cycles * interval  * (GetBlockIdx() - bigNum));
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + bigLength * interval * bigNum + smallLength * interval * (GetBlockIdx() - bigNum));
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + bigLength * cycles * interval  * bigNum + smallLength * cycles * interval  * (GetBlockIdx() - bigNum));
        }
        this->cycles = cycles;
        this->interval = interval;
        this->processDataLenght = interval;
        this->loopCount = coreDataNum;
        pipeIn->InitBuffer(inQueueInput, BUFFER_NUM, (interval) * sizeof(DTYPE_INPUT));
        pipeIn->InitBuffer(inQueueOther, BUFFER_NUM, (interval) * sizeof(DTYPE_INPUT));
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, (interval) * sizeof(DTYPE_INPUT));
    }
    __aicore__ inline void Process() {
        for (int32_t i = 0; i < this->loopCount; i++) {
            LocalTensor<DTYPE_INPUT> otherLocal = inQueueOther.AllocTensor<DTYPE_INPUT>();
            DataCopy(otherLocal, otherGm[i * this->processDataLenght], this->processDataLenght);
            inQueueOther.EnQue(otherLocal);
            otherLocal = inQueueOther.DeQue<DTYPE_INPUT>();
            for (int32_t j = 0; j < this->cycles; j++) {
                CopyIn(i*this->cycles + j, i);
                Compute(i, otherLocal);
                CopyOut(i*this->cycles + j);
            }
            inQueueOther.FreeTensor(otherLocal);
        }
    }
private:
    __aicore__ inline void CopyIn(int32_t progress, int32_t other_progress) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.AllocTensor<DTYPE_INPUT>();
        DataCopy(inputLocal, inputGm[progress * this->processDataLenght], this->processDataLenght);
        inQueueInput.EnQue(inputLocal);
    }
    __aicore__ inline void Compute(int32_t progress, LocalTensor<DTYPE_INPUT> otherLocal) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.DeQue<DTYPE_INPUT>();
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.AllocTensor<DTYPE_INPUT>();
        for(int32_t i=0; i<this->processDataLenght; i++)
        {
            outLocal.SetValue(i, lcm(inputLocal.GetValue(i), otherLocal.GetValue(i)));
        }
        inQueueInput.FreeTensor(inputLocal);
        outQueueOut.EnQue<DTYPE_INPUT>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.DeQue<DTYPE_INPUT>();
        DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->processDataLenght * sizeof(DTYPE_INPUT)), 0, 0, 0}; 
        DataCopyPad(outGm[progress * this->processDataLenght], outLocal, copyParams);
        outQueueOut.FreeTensor(outLocal);
    }
    __aicore__ inline DTYPE_INPUT lcm(DTYPE_INPUT m, DTYPE_INPUT n)
    {
        int64_t m1 = m;
        int64_t n1 = n;
        while(n1 != 0)
        {
            DTYPE_INPUT temp = m1 % n1;
            m1 = n1;
            n1 = temp;
        }
        int64_t mn = m / m1 * n;
        if(mn < 0) mn = -mn;
        return (DTYPE_INPUT)(mn);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInput, inQueueOther;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmpMask;
    GlobalTensor<DTYPE_INPUT> inputGm;
    GlobalTensor<DTYPE_INPUT> otherGm;
    GlobalTensor<DTYPE_INPUT> outGm;
    uint32_t loopCount;
    uint32_t cycles;
    uint32_t interval;
    uint32_t processDataLenght;
};
    constexpr uint64_t pre[65] = {0, 1ull,3ull,5ull,11ull,17ull,39ull,65ull,139ull,261ull,531ull,1025ull,2095ull,4097ull,8259ull,16405ull,32907ull,65537ull,131367ull,262145ull,524827ull,1048645ull,2098179ull,4194305ull,8390831ull,16777233ull,33558531ull,67109125ull,134225995ull,268435457ull,536887863ull,1073741825ull,2147516555ull,4294968325ull,8590000131ull,17179869265ull,34359871791ull,68719476737ull,137439215619ull,274877911045ull,549756338843ull,1099511627777ull,2199024312423ull,4398046511105ull,8796095120395ull,17592186061077ull,35184376283139ull,70368744177665ull,140737496778927ull,281474976710721ull,562949970199059ull,1125899906908165ull,2251799847243787ull,4503599627370497ull,9007199321981223ull,18014398509483025ull,36028797153190091ull,72057594038190085ull,144115188344291331ull,288230376151711745ull,576460752840837695ull,1152921504606846977ull,2305843010287435779ull,4611686018428436805ull,9223372039002292363ull};
class KernelLcmBC_1 {
public:
    __aicore__ inline KernelLcmBC_1() {}
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR other, GM_ADDR out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, 
                                uint32_t repeat, uint32_t cycles, uint32_t interval, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        uint32_t group = GetBlockIdx()/repeat;
        uint32_t within = GetBlockIdx()%repeat;
        if (within < bigNum) {
            coreDataNum = bigLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + group * cycles * interval + within * bigLength * interval);
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + group * interval);
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + group * cycles * interval + within * bigLength * interval);
        } else {
            coreDataNum = smallLength;
            inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input + group * cycles * interval + bigLength * interval * bigNum + smallLength * interval  * (within - bigNum));
            otherGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)other + group * interval);
            outGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)out + group * cycles * interval + bigLength * interval * bigNum + smallLength * interval  * (within - bigNum));
        }
        this->cycles = cycles;
        this->interval = interval;
        this->processDataLenght = interval;
        this->loopCount = coreDataNum;
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, interval * sizeof(DTYPE_INPUT));
    }
    __aicore__ inline void Process() {
        for (int32_t i = 0; i < this->loopCount; i++) {
            Compute(i);
            CopyOut(i);
        }
    }
private:
    __aicore__ inline void Compute(int32_t progress){
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.AllocTensor<DTYPE_INPUT>();
        for(int32_t i=0; i<this->processDataLenght; i+=8)
        {
            outLocal.SetValue(i, lcm(inputGm.GetValue(progress * this->interval + i), otherGm.GetValue(i)));
            outLocal.SetValue(i+1, lcm(inputGm.GetValue(progress * this->interval + i + 1), otherGm.GetValue(i + 1)));
            outLocal.SetValue(i+2, lcm(inputGm.GetValue(progress * this->interval + i + 2), otherGm.GetValue(i + 2)));
            outLocal.SetValue(i+3, lcm(inputGm.GetValue(progress * this->interval + i + 3), otherGm.GetValue(i + 3)));
            outLocal.SetValue(i+4, lcm(inputGm.GetValue(progress * this->interval + i + 4), otherGm.GetValue(i + 4)));
            outLocal.SetValue(i+5, lcm(inputGm.GetValue(progress * this->interval + i + 5), otherGm.GetValue(i + 5)));
            outLocal.SetValue(i+6, lcm(inputGm.GetValue(progress * this->interval + i + 6), otherGm.GetValue(i + 6)));
            outLocal.SetValue(i+7, lcm(inputGm.GetValue(progress * this->interval + i + 7), otherGm.GetValue(i + 7)));
        }
        outQueueOut.EnQue<DTYPE_INPUT>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<DTYPE_INPUT> outLocal = outQueueOut.DeQue<DTYPE_INPUT>();
        DataCopy(outGm[progress * this->interval], outLocal, this->processDataLenght);
        outQueueOut.FreeTensor(outLocal);
    }
    __aicore__ inline DTYPE_INPUT lcm(DTYPE_INPUT m, DTYPE_INPUT n)
    {
        DTYPE_INPUT a = m;
        DTYPE_INPUT b = n;
            DTYPE_INPUT shift = ScalarGetSFFValue<1>(a | b);
            a >>= ScalarGetSFFValue<1>(a);
            b >>= ScalarGetSFFValue<1>(b);
            while(b != 0) 
            {
                b >>= ScalarGetSFFValue<1>(b);
                if (a > b) {
                    a ^= b ^= a ^= b;
                }
                b -= a;
            }
            DTYPE_INPUT mn = (m/(a << shift))*n;
            mn = (mn > 0 ? mn : -mn);
            return mn;
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInput, inQueueOther;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmp5;
    GlobalTensor<DTYPE_INPUT> inputGm;
    GlobalTensor<DTYPE_INPUT> otherGm;
    GlobalTensor<DTYPE_INPUT> outGm;
    uint32_t loopCount;
    uint32_t cycles;
    uint32_t interval;
    uint32_t processDataLenght;
};
extern "C" __global__ __aicore__ void lcm(GM_ADDR input, GM_ADDR other, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) {
    TPipe pipe;
    if(TILING_KEY_IS(1)) {
        GET_TILING_DATA_WITH_STRUCT(LcmTilingData, tiling_data, tiling);
        KernelLcm op;
        op.Init(input, other, out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, &pipe);
        op.Process();
    }
    else if (TILING_KEY_IS(2)) {
        GET_TILING_DATA_WITH_STRUCT(LcmTilingDataBC, tiling_data, tiling);
        KernelLcmBC op;
        op.Init(input, other, out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, 
                tiling_data.repeat, tiling_data.cycles, tiling_data.interval, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(3)) {
        GET_TILING_DATA_WITH_STRUCT(LcmTilingDataBC, tiling_data, tiling);
        KernelLcmBC_1 op;
        op.Init(input, other, out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, 
                tiling_data.repeat, tiling_data.cycles, tiling_data.interval, &pipe);
        op.Process();
    }
}